
#include <fcntl.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
#include <unistd.h>
#include <sys/types.h>
#include <sys/stat.h>
#include <CL/cl.h>
#include <jni.h> 

const char *KernelSource = "\n" \
"__kernel void crypt(                                                    \n" \
"   __global char* text1, 	                                 			 \n" \
"   __global char* text2,                  				             	 \n" \
"   __global int* key,			     		                             \n" \
"   const unsigned int nthreads,                                       	 \n" \
"   const unsigned int array_rows                                      	 \n" \
")                                           						   	 \n" \
"{                                                                     	 \n" \
"	int id = get_global_id(0);											 \n" \
"if(id<nthreads){														 \n" \
"	int tslice = array_rows / 8;										 \n" \
"	int ttslice = (tslice + nthreads-1) / nthreads;						 \n" \
"	int slice = ttslice*8;												 \n" \
"	int ilow = id*slice;												 \n" \
"	int iupper = (id+1)*slice;										 	 \n" \
"	if(iupper > array_rows) iupper = array_rows;						 \n" \
"																		 \n" \
"	int i1 = ilow;														 \n" \
"	int i2 = ilow;														 \n" \
"	int ik;																 \n" \
"	unsigned int x1;													 \n" \
"	unsigned int x2;													 \n" \
"	unsigned int x3;													 \n" \
"	unsigned int x4;													 \n" \
"	int t1;																 \n" \
"	int t2;																 \n" \
"	int r;																 \n" \
"	int i;																 \n" \
"	for (i=ilow ; i <iupper ; i +=8)									 \n" \
"	{																	 \n" \
"	    ik = 0;															 \n" \
"	    r = 8;															 \n" \
"		x1 = text1[i1++] & 0xff;										 \n" \
"	    x1 |= (text1[i1++] & 0xff) << 8;								 \n" \
"	    x2 = text1[i1++] & 0xff;										 \n" \
"	    x2 |= (text1[i1++] & 0xff) << 8;								 \n" \
"	    x3 = text1[i1++] & 0xff;										 \n" \
"	    x3 |= (text1[i1++] & 0xff) << 8;								 \n" \
"	    x4 = text1[i1++] & 0xff;										 \n" \
"	    x4 |= (text1[i1++] & 0xff) << 8;								 \n" \
"																		 \n" \
"	    do {															 \n" \
"	        x1 = (int) ((long) x1 * key[ik++] % 0x10001L & 0xffff);		 \n" \
"	        x2 = x2 + key[ik++] & 0xffff;								 \n" \
"			x3 = x3 + key[ik++] & 0xffff;								 \n" \
"			x4 = (int) ((long) x4 * key[ik++] % 0x10001L & 0xffff);		 \n" \
"			t2 = x1 ^ x3;												 \n" \
"			t2 = (int) ((long) t2 * key[ik++] % 0x10001L & 0xffff);		 \n" \
"			t1 = t2 + (x2 ^ x4) & 0xffff;								 \n" \
"			t1 = (int) ((long) t1 * key[ik++] % 0x10001L & 0xffff);		 \n" \
"			t2 = t1 + t2 & 0xffff;										 \n" \
"			x1 ^= t1;													 \n" \
"			x4 ^= t2;													 \n" \
"			t2 ^= x2;													 \n" \
"			x2 = x3 ^ t1;												 \n" \
"			x3 = t2;													 \n" \
"	    } while(--r != 0);												 \n" \
"	   																  	 \n" \
"	    x1 = (int) ((long) x1 * key[ik++]  % (0x10001L) & 0xffff);		 \n" \
"	    x3 = x3 + key[ik++] & 0xffff;									 \n" \
"	    x2 = x2 + key[ik++] & 0xffff;									 \n" \
"	    x4 = (int) ((long) x4 * key[ik++]  % (0x10001L) & 0xffff);		 \n" \
"																		 \n" \
"	    text2[i2++] = (char) x1;										 \n" \
"	    text2[i2++] = (char) (x1 >> 8);							 		 \n" \
"	    text2[i2++] = (char) x3;										 \n" \
"	    text2[i2++] = (char) (x3 >> 8);									 \n" \
"	    text2[i2++] = (char) x2;										 \n" \
"	    text2[i2++] = (char) (x2 >> 8);									 \n" \
"	    text2[i2++] = (char) x4;										 \n" \
"	    text2[i2++] = (char) (x4 >> 8);									 \n" \
"	}																	 \n" \
"}																		 \n" \
"}																		 \n" \
"																		 \n";



////////////////////////////////////////////////////////////////////////////////
JNIEXPORT jbyteArray JNICALL Java_section2_crypt_IDEAOCLTest_cryptFromC
  (JNIEnv *env, jobject jobj, jbyteArray jplain, jintArray jZ, jintArray jDK, jint jnthreads, jint jarray_rows)
{
    int err;                            // error code returned from api calls
     
    jbyteArray res;          			// results returned from device
   
    size_t global,local;                      // global domain size for our calculation
 
    cl_device_id device_id;             // compute device id 
    cl_platform_id platform;		// compute platform id
    cl_context context;                 // compute context
    cl_command_queue commands;          // compute command queue
    cl_program program;                 // compute program
    cl_kernel kernel;                   // compute kernel


    struct timeval tim;
    double starttime,endtime,temp;

    cl_mem text1b;
    cl_mem text2b;
    cl_mem texttb;
    cl_mem keyb;
    cl_mem key2b;

   
    unsigned int nthreads = (unsigned int)jnthreads;
    unsigned int array_rows = (unsigned int)jarray_rows;
	
    char* plain = (*env)->GetByteArrayElements(env, jplain,0);
    int* Z = (*env)->GetIntArrayElements(env, jZ,0);
    int* DK = (*env)->GetIntArrayElements(env, jDK,0);


	
    res = (*env)->NewByteArray(env, array_rows);
    char* cryptres = malloc(array_rows * sizeof(char));
    char* result = malloc(array_rows * sizeof(char));

    /*//measure the opencl initialization
    gettimeofday(&tim, NULL); 
    starttime =  tim.tv_sec*1000+(tim.tv_usec*0.001);
*/

    // Connect to a compute device
   
    err = clGetPlatformIDs(1, &platform, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to select paltform!\n");
        return res;
    }/*
    gettimeofday(&tim, NULL); 
    endtime  = tim.tv_sec*1000+(tim.tv_usec*0.001);
    temp = (endtime-starttime)*0.001;
    printf("%4.3f ", temp);
*/
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to select device!\n");
        return res;
    }


    //
    context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
    if (!context)
    {
        printf("Error: Failed to create a compute context!\n");
        return res;
    }

    // Create a command queue
    //
    commands = clCreateCommandQueue(context, device_id, 0, &err);
    if (!commands)
    {
        printf("Error: Failed to create a command queue!\n");
        return res;
    }

    // Create the compute program from the source buffer
    //
    program = clCreateProgramWithSource(context, 1, (const char **) & KernelSource, NULL, &err);
    if (!program)
    {
        printf("Error: Failed to create compute program!\n");
        return res;
    }
    char* options = "-cl-denorms-are-zero";
    // Build the program executable
    //
    err = clBuildProgram(program, 0, NULL, options, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        size_t len;
        char buffer[2048];
 
        printf("Error: Failed to build program executable!\n");
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
        printf("%s\n", buffer);
        exit(1);
    }

    // Create the compute kernel in the program we wish to run
    //
    kernel = clCreateKernel(program, "crypt", &err);
    if (!kernel || err != CL_SUCCESS)
    {
        printf("Error: Failed to create compute kernel!\n");
        exit(1);
    }

   // Create arrays in device memory for our calculations
    //
    text1b = clCreateBuffer(context,  CL_MEM_READ_ONLY,  sizeof(char) * array_rows, NULL, NULL);
    text2b = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(char) * array_rows, NULL, NULL);
    keyb = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int) * 52, NULL, NULL);
	texttb = clCreateBuffer(context,  CL_MEM_READ_ONLY,  sizeof(char) * array_rows, NULL, NULL);
	key2b = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int) * 52, NULL, NULL);
    if (!text1b || !text2b || !keyb)
    {
        printf("Error: Failed to allocate device memory!\n");
        exit(1);       
    }    
    
   

   //ENCRYPT
    
    // Write our data sets into the arrays in device memory 
    //
    err = clEnqueueWriteBuffer(commands, text1b, CL_TRUE, 0, sizeof(char) * array_rows, plain, 0, NULL, NULL);
    err |= clEnqueueWriteBuffer(commands, keyb, CL_TRUE, 0, sizeof(int) * 52, Z, 0, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to write array to device memory!\n");
        exit(1);
    }

    // Set the arguments to our compute kernel
    //
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &text1b);
    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &text2b);
    err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &keyb);
    err |= clSetKernelArg(kernel, 3, sizeof(unsigned int), &nthreads);
    err |= clSetKernelArg(kernel, 4, sizeof(unsigned int), &array_rows);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to set kernel arguments! %d\n", err);
        exit(1);
    }

//get the number of compute units (SMs) of this device
    cl_uint units;
    clGetDeviceInfo( device_id, CL_DEVICE_MAX_COMPUTE_UNITS,  sizeof(cl_uint),  &units, NULL);

//get the multiple of local work group
    size_t mult;
    err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(mult), &mult, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to retrieve kernel work group info! %d\n", err);
        exit(1);
    }
  

    global = nthreads;
  
    local = global / (units*2); // create at least 2 warps per multiprocessor, so the SM can execute the other warp while the first one waits (a group has at least one warp)
    if(local>mult) {
    local = local / mult; //now make local a multiple of mult (warp size) so that resources in the warp are not wasted
    local = local * mult;
    }

    
    if(global % local != 0)
    {
		global=global-(global%local);
		global+=local; //make global a multiple of local
    }


    //measure 1st kernel exec time
    gettimeofday(&tim, NULL); 
    starttime =  tim.tv_sec*1000+(tim.tv_usec*0.001);

    // Wait for the commands to get serviced before reading back results
    //

    err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
    if (err!= CL_SUCCESS)
    {
        printf("Error: Failed to execute kernel!\n");
        return res;
    }

    clFinish(commands);
	
    gettimeofday(&tim, NULL); 
    endtime  = tim.tv_sec*1000+(tim.tv_usec*0.001);
    temp = (endtime-starttime)*0.001;
    

    // Read back the res from the device
    //
    err = clEnqueueReadBuffer( commands, text2b, CL_TRUE, 0, sizeof(char) * (array_rows), cryptres, 0, NULL, NULL );  
    
    if (err != CL_SUCCESS)
    {
        printf("-0.1");
        exit(1);
    }
    
   //DECRYPT
    // Write our data sets into the arrays in device memory 
    //
    err = clEnqueueWriteBuffer(commands, texttb, CL_TRUE, 0, sizeof(char) * array_rows, cryptres, 0, NULL, NULL);
    err |= clEnqueueWriteBuffer(commands, key2b, CL_TRUE, 0, sizeof(int) * 52, DK, 0, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to write array to device memory!\n");
        exit(1);
    }
    
    // Set the arguments to our compute kernel
    //
    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &texttb);
    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &text2b);
    err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &key2b);
    err |= clSetKernelArg(kernel, 3, sizeof(unsigned int), &nthreads);
    err |= clSetKernelArg(kernel, 4, sizeof(unsigned int), &array_rows);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to set kernel arguments! %d\n", err);
        exit(1);
    }
    //measure 2nd kernel exec time
    gettimeofday(&tim, NULL); 
    starttime =  tim.tv_sec*1000+(tim.tv_usec*0.001);
    err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
    if (err)
    {
        printf("Error: Failed to execute kernel!\n");
        return res;
    }
   
    // Wait for the commands to get serviced before reading back results
    //
    clFinish(commands);
   
    gettimeofday(&tim, NULL); 
    endtime  = tim.tv_sec*1000+(tim.tv_usec*0.001);
    temp += (endtime-starttime)*0.001;
    printf("%4.3f", temp);

    // Read back the res from the device
    //
    err = clEnqueueReadBuffer( commands, text2b, CL_TRUE, 0, sizeof(char) * (array_rows), result, 0, NULL, NULL );  
    
    if (err != CL_SUCCESS)
    {
        printf("error");
        exit(1);
    }

    // Shutdown and cleanup
    //
    (*env)->SetByteArrayRegion(env,res, 0, array_rows, result);

	free(cryptres);
	free(result);
    clReleaseMemObject(text1b);
    clReleaseMemObject(text2b);
    clReleaseMemObject(keyb);
    clReleaseProgram(program);
    clReleaseKernel(kernel);
    clReleaseCommandQueue(commands);
    clReleaseContext(context);

	

    
    return res;
}




 
    
